__kernel void gpuPdist (
	__global float *in,
	__global float *out,
	__global int n,
	__global int m,
	__local float * Ys, //[16][16] pass in 2d arrays
	__local float * Xs ) //[16][16]
{ 
	int bx = get_group_id(0);
	int by = get_group_id(1);//int bx = blockIdx.x, by = blockIdx.y;
	int tx = get_local_id(0);
	int ty = get_local_id(1);//int tx = threadIdx.x, ty = threadIdx.y;
	int yBegin = by * 16 * m;
	int xBegin = bx * 16 * m;
	int yEnd = yBegin + m - 1;
	int y, x, k, o;
	float s = 0;
	float tmp;
  
    x = xBegin;
	for(y = yBegin; y <= yEnd ; y += 16)
	{ 
		Ys[tx * 16 + ty] = in[ty * m + y + tx];
		Xs[ty * 16 + tx] = in[ty * m + x + tx];
		 //*** note the transpose of Xs 
		barrier(CLK_LOCAL_MEM_FENCE);
		for(k = 0; k < 16; k++)
		{ 
			tmp = Ys[k * 16 + ty] - Xs[tx * 16 + k];
			s += tmp*tmp;
		}
		x += 16;
		barrier(CLK_LOCAL_MEM_FENCE);
	}
	o = by*16*n + ty*n + bx*16 + tx;
	out[o] = sqrt(s);
	barrier(CLK_LOCAL_MEM_FENCE);
}